In `@examples/distributed/intranode/example_alltoall_route2x4.py`:
- Line 42: The kernel signature currently includes an unused parameter barrier;
update the kernel or its invocation to remove the unused parameter or use it for
synchronization: either delete barrier from the kernel signature and also stop
allocating/passing it in run_torus_alltoall, or add explicit barrier
synchronization inside the kernel (referencing the kernel function where
barrier: T.Tensor((PE_num), "int32") is declared) and ensure run_torus_alltoall
continues to allocate and pass the barrier tensor so the symbol names match
(update both the kernel definition and the call sites in run_torus_alltoall).
- Around line 279-288: The script does not validate that X * Y equals PE_num
before spawning processes, which can cause incorrect torus topology; after
parsing args in the main block (the argparse code that sets args.M, args.N,
args.PE_num, args.X, args.Y) add a check that args.X * args.Y == args.PE_num and
raise a clear error (ValueError or argparse.ArgumentTypeError) or call
parser.error with a descriptive message if the equality fails so
run_torus_alltoall is only invoked with a consistent torus configuration.

In `@examples/distributed/intranode/example_alltoall.py`:
- Around line 73-75: The bandwidth formula is computing elements/sec but not
bytes/sec; update the print calculation to multiply by the element size (in
bytes) — e.g., use dtype_itemsize = np.dtype(tensor_dtype).itemsize or hard-code
2 for float16 — then compute total_bytes = 2 * PE_num * M * N * dtype_itemsize
and compute bandwidth_gbps = total_bytes / (elapsed_time * 1e6) (use the same
elapsed_time units as the printed ms) and print bandwidth_gbps; update the
expression where local_rank, PE_num, M, N, elapsed_time are used to include this
element-size factor.
- Around line 12-37: The kernel declares a barrier parameter that is never used
in main; either remove it from the kernel signature (and corresponding callers)
or implement inter-rank synchronization using that barrier instead of relying
solely on T.fence_sys. If you choose to use the barrier, after T.put_block and
T.fence_sys update the barrier tensor to signal this rank’s completion and then
wait/spin until all PE_num entries are set (i.e., implement a full all-ranks
barrier using the barrier buffer), referencing the symbols main, barrier,
T.put_block and T.fence_sys; if you choose removal, delete the barrier parameter
from main and all call sites.

In `@src/op/remote_copy.cc`:
- Around line 422-431: The code constructs a function name using sem_str[sem]
and scope_str[scope] without validating indices, which can lead to buffer
overreads; add bounds checks for sem (<6) and scope (<4) before indexing
(similar to StOpNode::Lower and LdOpNode::Lower), and handle out-of-range values
by selecting a safe default string or emitting an error/log and returning early
so the use in building "tl::ptx_atom_add_<sem>_<scope>" cannot index past
sem_str or scope_str.

In `@src/op/remote_copy.h`:
- Around line 327-328: Update the field comments for the "scope" and "sem"
members to match the actual enum/mapping used in the implementation: describe
scope as "{0: cta, 1: cluster, 2: gpu, 3: sys}" and describe sem as "{0: weak,
1: volatile, 2: relaxed, 3: acquire, 4: release, 5: acq_rel}" so the
documentation for the int members scope and sem in remote_copy.h matches the
implemented semantics.

In `@tilelang/language/builtin.py`:
- Around line 839-854: The is_remote check in the function incorrectly treats a
plain Python int dst_pe == -1 as remote because it only checks
isinstance(dst_pe, tir.IntImm); update the logic in the is_remote computation
(the variable named is_remote) to also consider a plain int -1 as local—e.g.,
treat dst_pe == -1 the same as tir.IntImm with value == -1 or normalize dst_pe
via getattr(dst_pe, "value", dst_pe) before comparison—so the remote branch (the
code that builds local_rank, local_base_ptr, offset_to_base, remote_ptr and
calls tir.call_extern with remote_ptr) is only taken when dst_pe truly indicates
a remote PE.
- Around line 737-738: Update the incorrect docstring on the function atom_add:
replace the copy-pasted PTX cp.async barrier description with a concise
description stating that atom_add performs an atomic add operation (e.g.,
performs an atomic addition of a value to a target and returns the previous
value or the new value per implementation), mention relevant semantics such as
data type(s) supported and whether the operation is relaxed/ordered or maps to a
specific PTX atomic instruction, and ensure the docstring matches the function
signature and behavior in atom_add.
- Around line 733-747: The atom_add function currently uses direct dict lookups
for scope_str and sem_str which will raise KeyError for unsupported MemoryScope
or MemorySemantic values; update atom_add to explicitly validate inputs before
using them (check the provided scope against allowed values {MemoryScope.GPU,
MemoryScope.SYSTEM} and sem against {MemorySemantic.RELAXED,
MemorySemantic.ACQUIRE, MemorySemantic.RELEASE, MemorySemantic.ACQ_REL}), and if
invalid raise a clear ValueError mentioning the parameter name, the unsupported
value, and the list of supported options; then proceed to build
scope_str/sem_str and call tir.call_intrin(address_of(barrier), value, sem_str,
scope_str) as before.

In `@examples/distributed/intranode/example_alltoall_route2x4.py`:
- Around line 325-327: The three print statements use f-strings without any
placeholders (e.g., print(f"Benchmark Results:"), print(f"  Average Latency:
{elapsed_time_ms:.4f} ms") — actually the latter two do use placeholders so only
remove the redundant f-prefix from lines that have no placeholders; specifically
change print(f"Benchmark Results:") to print("Benchmark Results:"), and likewise
any other literal-only prints in the same block, leaving prints that format
variables (like elapsed_time_ms and bandwidth_gbps) as f-strings; update the
print statements in the block containing "Benchmark Results:", "  Average
Latency:", and "  Effective Bandwidth:" accordingly.
- Around line 233-236: Validate and guard block sizing before launching the
kernel: ensure block_M is > 0 and that M is evenly divisible by block_M (or
implement tail handling) so the kernel’s copy of block_M * N elements cannot go
out-of-bounds; specifically, check/adjust the variables M, block_M, block_N and
threads used in the setup (where block_M = M // 8, block_N = N, num_blocks_M = M
// block_M) to either compute num_blocks_M = ceil_div(M, block_M) and handle the
final partial block inside the kernel or error out with a clear message when
block_M == 0 or M % block_M != 0.
- Around line 27-28: The function torus_alltoall_xy currently declares a
parameter block_N which is never used; remove block_N from the function
signature and from all call-sites that pass it (ensure callers now pass only
PE_num, X, Y, M, N, block_M, threads), and update any related documentation or
tests that reference the old signature; search for usages of torus_alltoall_xy
to update invocations and adjust any higher-level wrappers or type hints
referencing block_N to keep the API consistent.

In `@examples/distributed/intranode/example_alltoall_route2x4.py`:
- Around line 303-316: The benchmark loops reuse tensors so local_finish,
global_finish, and signal_transfer retain values across runs causing kernel (the
all-to-all transfer loop) to skip work; before each warmup and each timed
iteration (i.e., before calling kernel in both the 5 warmup iterations and
inside the num_iters loop) reset/zero-out those tensors (local_finish,
global_finish, signal_transfer) on the appropriate device so the kernel’s
while-loop condition behaves correctly; locate references to kernel(...,
buffer_transfer, signal_transfer, local_finish, global_finish, barrier) and
insert resets just prior to those calls.

In `@examples/distributed/intranode/example_alltoall_route2x4.py`:
- Around line 236-243: The host-side block computation can produce num_blocks ==
0 and trigger a ZeroDivisionError when computing tile_M; modify the calculation
of num_blocks (used before computing tile_M) to guard against zero by ensuring
num_blocks = max(1, min(num_tiles, 148 // (PE_num * PE_num))) (or explicitly
check if num_blocks == 0 and set it to 1) so the subsequent tile_M = M //
(num_blocks * num_warps) cannot divide by zero; update references to num_blocks,
num_tiles, num_warps and tile_M accordingly.
- Around line 27-35: In torus_alltoall_xy, remove the hardcoded num_SM=148 and
either query the device SM count (e.g., cuda.get_device().MULTIPROCESSOR_COUNT
or similar) or add num_SM as a function parameter; remove or use the unused
tile_N parameter (Ruff ARG001) — if block_N was intended, rename consistently;
compute num_blocks = max(1, min(num_SM // (PE_num * PE_num), num_tiles)) to
avoid ZeroDivisionError when PE_num is large, then compute block_M,
tiles_per_block and keep the existing tiles_per_block <= num_warps assertion
using num_warps = threads // 32. Ensure all references use the updated names
(num_SM, num_blocks, block_M, tiles_per_block, num_warps).

In `@examples/distributed/intranode/example_alltoall_route2x4.py`:
- Around line 20-26: The decorator call to tilelang.jit currently passes a
hardcoded developer-specific debug_root_path
(debug_root_path="/home/zhengju.tang/...") which must be removed or made
portable; update the tilelang.jit invocation (the pass_configs and debug
options) to either omit the debug_root_path argument entirely or replace it with
a relative path or environment-driven value (e.g., read from an env var or use
"./debug") so the example contains no absolute user home directories.
- Around line 114-157: chunk_size can become negative when tile_M < PE_num;
before calling T.put_warp (both branches) compute and guard chunk_size so it is
non-negative—for example clamp available = T.max(0, tile_M - chunk_start) and
set chunk_size = T.min(chunk_M, available) (or skip the
T.put_warp/T.atom_add/T.sync_warp sequence when chunk_size <= 0). Update the
calculations around chunk_M, chunk_start, chunk_size and the put_warp calls to
use this non-negative-clamped value (references: chunk_M, chunk_start,
chunk_size, bx, tile_M, PE_num, and the T.put_warp/T.atom_add calls).

In `@examples/distributed/intranode/example_alltoall_route2x4.py`:
- Around line 165-166: The condition uses a boolean comparison "flag[0] ==
False"; change it to use a truthiness check by replacing "flag[0] == False" with
"not flag[0]" in the if statement that checks signal_transfer, dst_rank and flag
(the line referencing signal_transfer[bx, dst_rank[0], bz, warp_idx] == PE_num
and flag[0] == False) so it conforms to PEP8 and handles non-boolean flag values
correctly.
- Line 245: The host-side assignment tile_M = M // (num_blocks * num_warps) is
dead code (unused) because the kernel computes its own tile_M; remove this
unused variable assignment (or instead pass it into the kernel if you intended
to use it) to fix the Ruff F841 warning—look for the symbol tile_M and the
related variables M, num_blocks, num_warps in example_alltoall_route2x4.py and
either delete the assignment or wire tile_M into the kernel invocation so it is
actually consumed.

In `@examples/distributed/intranode/example_alltoall_route2x4.py`:
- Around line 276-278: The shape args passed to tilelang.tensor are plain
integers because (1) and (PE_num) are not tuples; change the single-element
shapes to proper 1-tuples when creating tensors in local_finish, global_finish
and barrier by passing (1,) and (PE_num,) to tilelang.tensor (keep allocator and
types the same).